#pragma once

#if defined(__cplusplus) && defined(__CUDACC__)

#include "builtin_types.h"
#include "device_types.h"
#include "host_defines.h"

#if defined(__CUDACC_RTC__)
#define __DEVICE_FUNCTIONS_DECL__ __device__ __cudart_builtin__
#define __DEVICE_FUNCTIONS_STATIC_DECL__ __device__ __cudart_builtin__
#else
#define __DEVICE_FUNCTIONS_DECL__ __device__ __cudart_builtin__
#define __DEVICE_FUNCTIONS_STATIC_DECL__ static __inline__ __device__ __cudart_builtin__
#endif /* __CUDACC_RTC__ */

extern "C"
{
    /* TODO: Those functions are not implemented by compiler or kernel lib yet. */
    __DEVICE_FUNCTIONS_DECL__ __device_builtin__ unsigned int __pm0(void);
    __DEVICE_FUNCTIONS_DECL__ __device_builtin__ unsigned int __pm1(void);
    __DEVICE_FUNCTIONS_DECL__ __device_builtin__ unsigned int __pm2(void);
    __DEVICE_FUNCTIONS_DECL__ __device_builtin__ unsigned int __pm3(void);
    __DEVICE_FUNCTIONS_DECL__ __device_builtin__ void __prof_trigger(int);
}

#if defined(_WIN32)
#define __DEVICE_FUNCTIONS_DEPRECATED__(msg) __declspec(deprecated(msg))
#elif (defined(__GNUC__) && (__GNUC__ < 4 || (__GNUC__ == 4 && __GNUC_MINOR__ < 5 && !defined(__clang__))))
#define __DEVICE_FUNCTIONS_DEPRECATED__(msg) __attribute__((deprecated))
#else
#define __DEVICE_FUNCTIONS_DEPRECATED__(msg) __attribute__((deprecated(msg)))
#endif

#define ___DEVICE_FUNCTIONS_STRINGIFY_INNERMOST(x) #x
#define __DEVICE_FUNCTIONS_STRINGIFY(x) ___DEVICE_FUNCTIONS_STRINGIFY_INNERMOST(x)

#define __DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(x) \
    __DEVICE_FUNCTIONS_STRINGIFY(x)                \
    "() is deprecated in favor of __" __DEVICE_FUNCTIONS_STRINGIFY(x) "() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning)."
#define __DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(x, y) \
    __DEVICE_FUNCTIONS_STRINGIFY(x)                   \
    "() is deprecated in favor of __" __DEVICE_FUNCTIONS_STRINGIFY(x) __DEVICE_FUNCTIONS_STRINGIFY(y) "() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning)."

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mulhi)) int mulhi(const int a, const int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mulhi)) unsigned int mulhi(const unsigned int a, const unsigned int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mulhi)) unsigned int mulhi(const int a, const unsigned int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mulhi)) unsigned int mulhi(const unsigned int a, const int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mul64hi)) long long int mul64hi(const long long int a, const long long int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mul64hi)) unsigned long long int mul64hi(const unsigned long long int a, const unsigned long long int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mul64hi)) unsigned long long int mul64hi(const long long int a, const unsigned long long int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mul64hi)) unsigned long long int mul64hi(const unsigned long long int a, const long long int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(float_as_int)) int float_as_int(const float a);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(int_as_float)) float int_as_float(const int a);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(float_as_uint)) unsigned int float_as_uint(const float a);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(uint_as_float)) float uint_as_float(const unsigned int a);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(saturate, f)) float saturate(const float a);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(mul24)) int mul24(const int a, const int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1(umul24)) unsigned int umul24(const unsigned int a, const unsigned int b);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(float2int, _ru | _rd | _rn | _rz)) int float2int(const float a, const enum cudaRoundMode mode = cudaRoundZero);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(float2uint, _ru | _rd | _rn | _rz)) unsigned int float2uint(const float a, const enum cudaRoundMode mode = cudaRoundZero);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(int2float, _ru | _rd | _rn | _rz)) float int2float(const int a, const enum cudaRoundMode mode = cudaRoundNearest);

__DEVICE_FUNCTIONS_STATIC_DECL__ __DEVICE_FUNCTIONS_DEPRECATED__(__DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2(uint2float, _ru | _rd | _rn | _rz)) float uint2float(const unsigned int a, const enum cudaRoundMode mode = cudaRoundNearest);
#undef __DEVICE_FUNCTIONS_DEPRECATED__
#undef ___DEVICE_FUNCTIONS_STRINGIFY_INNERMOST
#undef __DEVICE_FUNCTIONS_STRINGIFY
#undef __DEVICE_FUNCTIONS_DEPRECATION_MESSAGE1
#undef __DEVICE_FUNCTIONS_DEPRECATION_MESSAGE2

#undef __DEVICE_FUNCTIONS_DECL__
#undef __DEVICE_FUNCTIONS_STATIC_DECL__

#endif /* __cplusplus && __CUDACC__ */
#if !defined(__CUDACC_RTC__)
#include "device_functions.hpp"
#endif /* !defined(__CUDACC_RTC__) */
// TODO: Those files are not implemented yet.
// #include "device_atomic_functions.h"
#include "device_double_functions.h"
// #include "sm_20_atomic_functions.h"
// #include "sm_32_atomic_functions.h"
// #include "sm_35_atomic_functions.h"
// #include "sm_60_atomic_functions.h"
// #include "sm_20_intrinsics.h"
// #include "sm_30_intrinsics.h"
// #include "sm_32_intrinsics.h"
// #include "sm_35_intrinsics.h"
// #include "sm_61_intrinsics.h"
#include "sm_70_rt.h"
#include "sm_80_rt.h"
#include "surface_functions.h"
#include "texture_fetch_functions.h"
#include "texture_indirect_functions.h"
#include "surface_indirect_functions.h"

